#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <cuda_runtime.h>
#include <cutil_inline.h>

int numberOfBlocks = 10; // any number for now
int numberOfThreads = 256; // 
int n = 1048576;
int threadArraySize ;
bool* data;

// kernel code
__global__ void myFirstKernel(bool *DeviceArr, int globalVariable, int threadArraySize  )
{
	int start;
	int idx = blockIdx.x*blockDim.x + threadIdx.x;
	if (DeviceArr[globalVariable]) {
		start = /*(globalVariable*globalVariable) + */ (threadArraySize * idx /globalVariable  * globalVariable);
		for ( ; start <= threadArraySize *(idx+1); start += globalVariable)
			if (start!= globalVariable)
				DeviceArr[start] = false;
	}
}

void checkGpuMem(unsigned long long size)
{
	float free_m,total_m,used_m,mem_used,temp1,here1;
	size_t free_t,total_t, temp;
	unsigned int mem, rana,here;

	cudaMemGetInfo(&free_t,&total_t);	
	free_m =(unsigned int)free_t/1048576.0;
	total_m=(unsigned int)total_t/1048576.0;
	temp=(unsigned int)size/1048576.0;
	used_m=total_m-free_m;
	here1=free_t-(unsigned int )size/1048576.0;
	mem=free_m-temp;
	rana=temp-free_m;
	printf ( "mem free %d .... %f MB mem \ntotal %d....%f MB mem \nused %f MB\n",free_t,free_m,total_t,total_m,used_m);
	printf(" mem free after array %i MB\n",mem);
	printf(" negative mem free after array %i MB\n",rana);
	printf(" bytes mem free after array %i MB\n",here1);
}

int main () {

	unsigned long long size=n*1;
	unsigned int timer=0;
	checkGpuMem(size);
	//printf("%i memory used ",size);
	int i;
	bool* hostArr;
	bool* deviceArr;
	data = (bool*)malloc((n+1)*sizeof(bool));
	data[0] = false;
	data[1] = false;
	for(i=2; i<= n; i++)
		data[i]=true;
	hostArr = data ;
	size_t memSize = n*sizeof(bool);
	cudaMalloc( (void **) &deviceArr, memSize );
	cudaMemcpy( deviceArr, hostArr, memSize, cudaMemcpyHostToDevice );
	dim3 dimGrid( numberOfBlocks );
	dim3 dimBlock( numberOfThreads );
	cutilCheckError( cutCreateTimer( &timer));
	cutilCheckError( cutStartTimer( timer));
	cudaEvent_t start, stop;
	float time;
	cudaEventCreate(&start);
	cudaEventCreate(&stop);
	cudaEventRecord( start, 0 );
	threadArraySize = ceil((double)(n/*-i*i*/)/(numberOfThreads*numberOfBlocks)) ;
	for(int i=2; i<sqrt((double)n); i++)
	{
		myFirstKernel<<< dimGrid, dimBlock >>>(deviceArr, i, threadArraySize );
	}
	cutilCheckError( cutStopTimer( timer));
	cudaEventRecord( stop, 0 );
	cudaEventSynchronize( stop );
	cudaEventElapsedTime( &time, start, stop );
	cudaEventDestroy( start );
	cudaEventDestroy( stop );

	printf("Processing time: %f (ms)\n", cutGetTimerValue( timer));
	printf("\nTotal Time = %.5f seconds\n",time/1000);
	int primeCount = 1;
	cudaMemcpy( hostArr,deviceArr , memSize, cudaMemcpyDeviceToHost );
	for (int i = 3 ; i < n ; i +=2)
		if (hostArr[i]) {
			primeCount++;
	//		 printf("%i ", i);
		}
	//checkGpuMem();
	printf("%i\n",primeCount);
	cudaFree(deviceArr);
	free(hostArr);
	scanf("%i", n);
}
